ctran: enable per-block flag for broadcast with tcpdm/unpack#2079
Closed
fomichev wants to merge 2 commits intometa-pytorch:mainfrom
Closed
ctran: enable per-block flag for broadcast with tcpdm/unpack#2079fomichev wants to merge 2 commits intometa-pytorch:mainfrom
fomichev wants to merge 2 commits intometa-pytorch:mainfrom
Conversation
Summary: I know this is not something that's gonna be supported for long or even used in prod, but it's working with a minimal code changes, so re-enabled v2.29 for conda with iter builds. 901 + ncclx + tcpdm, conda package: 0f473b3 Differential Revision: D100339239
Summary: The TCPDM broadcast kernel (ncclKernelBroadcast<UNPACK=true>) hangs with multiple GPUs because all 8 CUDA thread blocks share a single kernel flag (flag[0]), creating a race condition on termination. Send/recv kernels don't have this problem because they use per-block flags (flag[blockIdx.x]): each block signals and terminates on its own flag slot, so block 0 clearing its slot doesn't affect other blocks. Relevant for TCPDM only because we have more than one block (due to unpack). Reviewed By: function47 Differential Revision: D99452874
meta-cla
bot
added
the
CLA Signed
Contributor
Contributor
|
This pull request has been merged in a5eae6f. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary:
The TCPDM broadcast kernel (ncclKernelBroadcast<UNPACK=true>) hangs
with multiple GPUs because all 8 CUDA thread blocks share a single
kernel flag (flag[0]), creating a race condition on termination.
Send/recv kernels don't have this problem because they use per-block
flags (flag[blockIdx.x]): each block signals and terminates on its
own flag slot, so block 0 clearing its slot doesn't affect other blocks.
Relevant for TCPDM only because we have more than one block (due to unpack).
Reviewed By: function47
Differential Revision: D99452874